#include <iostream>
#include <cuda_runtime.h>
#include <random>
#include <ctime>
#include <sys/time.h>

const int WARP_SIZE = 32;

const int BLOCK_SIZE = 256;

#define PERF(name) Perf perf_##name##__COUNTER__(#name)
#define PERF_CPU(name) PerfCPU perf_CPU_##name##__COUNTER__(#name)

class PerfCPU
{
public:
    PerfCPU(const std::string& name) {
        m_name = name;
        gettimeofday(&m_start, NULL);
    }

    ~PerfCPU() {
        gettimeofday(&m_end, NULL);
        float elapsed_time = (m_end.tv_sec - m_start.tv_sec) * 1000.0 + (m_end.tv_usec - m_start.tv_usec) / 1000.0;
        std::cout << m_name << " elapse: " << elapsed_time << " ms" << std::endl;
    }

private:
    std::string m_name;
    struct timeval m_start, m_end;
}; // class PerfCPU

class Perf
{
public:
    Perf(const std::string& name) {
        m_name = name;
        cudaEventCreate(&m_start);
        cudaEventCreate(&m_end);
        cudaEventRecord(m_start);
        cudaEventSynchronize(m_start);
    }

    ~Perf() {
        cudaEventRecord(m_end);
        cudaEventSynchronize(m_end);
        float elapsed_time = 0.0;
        cudaEventElapsedTime(&elapsed_time, m_start, m_end);
        std::cout << m_name << " elapse: " << elapsed_time << " ms" << std::endl;
    }

private:
    std::string m_name;
    cudaEvent_t m_start, m_end;
}; // class Perf


__global__ void sm_test_constant(float* input, float* output)
{
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    float tmp = input[idx];
    float out = 0.0;
    float theta = 10e-3;
    while (tmp > 10e-3) {
        out += 1.0;
        tmp -= 1.0;
    }
    output[idx] = out;
    return;
}

__global__ void sm_test_reg(float* input, float* output)
{
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    float tmp = input[idx];
    float out = 0.0;
    float theta = 10e-3;
    while (tmp > theta) {
        out += 1.0;
        tmp -= 1.0;
    }
    output[idx] = out;
    return;
}

int main(int argc, char* argv[])
{
    int block_dim = BLOCK_SIZE;
    int grid_dim = 1024000 / block_dim;
    
    // cpu mem
    float* input = new float[grid_dim * block_dim];
    float* output = new float[grid_dim * block_dim];

    std::default_random_engine generator;
    std::uniform_real_distribution<float> distribution(1.0, 10000.0);
    for (int i = 0; i < grid_dim * block_dim; ++i) {
        input[i] = 1000.0;
    }

    // gpu mem
    float* d_input;
    float* d_output;
    cudaMalloc((void**)&d_input, grid_dim * block_dim * sizeof(float));
    cudaMalloc((void**)&d_output, grid_dim * block_dim * sizeof(float));

    cudaMemcpy(d_input, input, grid_dim * block_dim * sizeof(float), cudaMemcpyHostToDevice);

    sm_test_constant<<<grid_dim, block_dim>>>(d_input, d_output);
    sm_test_reg<<<grid_dim, block_dim>>>(d_input, d_output);

    cudaMemcpy(output, d_output, grid_dim * block_dim * sizeof(float), cudaMemcpyDeviceToHost);
    cudaFree(d_input);
    cudaFree(d_output);

    delete[] input;
    delete[] output;
    return 0;
}